

#define CurrentID ((get_group_id(1)*get_num_groups(0)*get_num_groups(2)) + (get_group_id(2)*get_num_groups(0)) + get_group_id(0))
#define BlockNum 16

#define XDomainMax 65535
#define YDomainMax 65535
#define ZDomainMax 65535

uint f_BinarySearch(__global ulong2* Hash, uint domainNum, ulong var);

//通过哈希值取得坐标
inline ushort3 f_getHashCoord(const ulong hash){
    ushort3 coord;
    coord.x = hash;
    coord.y = (hash>>16);
    coord.z = (hash>>32);
    return coord;
}
inline ulong f_getCoordHash(const ushort3 coord){
    return upsample((uint)(coord.z), ((uint)coord.y<<16)|coord.x);
}

//采样取得六个轴向的坐标哈希值
inline ulong f_SamplingHash(ushort3 coord, uchar dir) {
    /*
     |
     |-----------| |-----------| |-----------|
     | . | . | . | | . | 4 | . | | . | . | . |
     |-----------| |-----------| |-----------|
     | . | 2 | . | | 1 | . | 0 | | . | 3 | . |
     |-----------| |-----------| |-----------|
     | . | . | . | | . | 5 | . | | . | . | . |
     |-------------------------------------------x
     */
    
    switch(dir){
        case 0:
            if(coord.x+1 < XDomainMax) return upsample((uint)(coord.z), upsample(coord.y, coord.x+1));
            break;
        case 1:
            if(coord.x) return upsample((uint)(coord.z), upsample(coord.y, coord.x-1));
            break;
        case 2:
            if(coord.y+1 < YDomainMax) return upsample((uint)(coord.z), upsample((ushort)(coord.y+1), coord.x));
            break;
        case 3:
            if(coord.y) return upsample((uint)(coord.z), upsample((ushort)(coord.y-1), coord.x));
            break;
        case 4:
            if(coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y, coord.x));
            break;
        case 5:
            if(coord.z) return upsample((uint)(coord.z-1), upsample(coord.y, coord.x));
            break;
        
    }
    
    return 0x1000000000000;
}

//通过哈希值找到值所在数组位置
uint f_BinarySearch(__global ulong2* Hash, uint domainNum, ulong var){
    //二分查找哈希
    uint HashIndex = domainNum/2;
    if (var<0x1000000000000) {
        ulong domain = 0;
        for (uint i = HashIndex; i < domainNum; ) {
            //if(HashIndex<32) break;
            domain = Hash[i].x;
            
            if (domain == var) {
                return i;
            } else if (HashIndex == 0) {
                break;
            } else {
                if(HashIndex<16){
                    HashIndex = i;
                    break;
                }
                HashIndex /= 2;
                if (domain < var) {
                    i += HashIndex;
                } else if (domain > var) {
                    i -= HashIndex;
                }
            }
        }
        for (uint i = 0; i < 16; ++i) {
            if(var == Hash[HashIndex+i].x){
                return HashIndex+i;
            }
        }
        for (uint i = 0; i < 16; ++i) {
            if(var == Hash[HashIndex-i].x){
                return HashIndex-i;
            }
        }
    }
    
    return 0xffffffff;
}


//记录有面片的哈希值
void f_RecordPlaceholder(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum);
void f_RedixSort_OffsetA(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uint offset);
void f_RedixSort_OffsetB(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uint offset);
void f_MeshAssemble(__global float4* vertex, __global float4* normal, __global ulong2* hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum);


//根据哈希值设置体素类型 设置重新拍下偏移
__kernel void f_ParticleMesh(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum){
	const size_t id = CurrentID;
	if(id<eNum){
		uint grid[6];
        uchar maskNum = 0;
        Hash[id].s1 = Hash[id].s0;
        uchar fceasNum = 6;
        //哈希值转坐标
        ushort3 coord = f_getHashCoord(Hash[id].s1);
		//uint localID =  get_local_id(0);
        for(uchar localID=0; localID<6; ++localID){
            //坐标去六轴哈希值
            uint grid = f_BinarySearch(Hash, eNum, f_SamplingHash(coord, localID));
            if(grid < 0xffffffff){
                maskNum |= (1<<localID);
                --fceasNum;
            }
        }

        //maskNum |= (1<<1);
        
        if(maskNum < 63){
            rangeNum[id] = maskNum;
            rangeOffset[id] = fceasNum*4;
        }else{
            rangeNum[id] = 0xff;
            rangeOffset[id] = 0;
        }
        
        
		vertex[id] = (float4)(10,1,1,1);
		if(id == 0){
			//去除完全包围的粒子
			size_t work_size[3];
			work_size[0] = get_num_groups(0);
			work_size[1] = get_num_groups(1);
			work_size[2] = get_num_groups(2);
			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetA(vertex, normal, Hash, rangeOffset, rangeNum, eNum, 1); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);

		}
	}
}


//=================================开始索引累加======================================
void f_RedixSort_OffsetA(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uint offset){
	const size_t id = CurrentID;
	size_t CurrentOffset = id+offset;
	if (CurrentOffset < eNum) rangeOffset[CurrentOffset].s0 += rangeOffset[id].s1;
		
	//barrier(CLK_GLOBAL_MEM_FENCE);
	//===============偏移拷贝======================
	if(id==0){
		size_t work_size[3];
		if(offset<=eNum){
			work_size[0] = pow(eNum-offset, 1.0f/3.0f)+1; 
			work_size[1] = work_size[0]; 
			work_size[2] = work_size[0];

			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetB(vertex, normal, Hash, rangeOffset, rangeNum, eNum, offset); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
			//if(err_ret<0) printf("Previous copy D  = %u \n", err_ret);
		} else {
			//拼装顶点
			work_size[0] = pow(eNum, 1.0f/3.0f)+1; 
			work_size[1] = work_size[0]; 
			work_size[2] = work_size[0]; 
			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_MeshAssemble(vertex, normal, Hash, rangeOffset, rangeNum, eNum); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
		}
	}
}
void f_RedixSort_OffsetB(__global float4* vertex, __global float4* normal, __global ulong2* Hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uint offset){
	const size_t id = CurrentID;
	size_t CurrentOffset = id+offset;
	
	if(CurrentOffset<eNum) rangeOffset[CurrentOffset].s1 = rangeOffset[CurrentOffset].s0;
	
	//barrier(CLK_GLOBAL_MEM_FENCE);
	//===============偏移===================
	if(id==0){
		size_t work_size[3];
		work_size[0] = get_num_groups(0);
		work_size[1] = get_num_groups(1);
		work_size[2] = get_num_groups(2);
		ndrange_t ndrange1 = ndrange_3D(work_size);
		void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetA(vertex, normal, Hash, rangeOffset, rangeNum, eNum, offset*2); };
		int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
		//if(err_ret<0) printf("Previous copy D  = %u \n", err_ret);
	}
}



/*
 
 0 +X
 1 -X
 2 +Y
 3 -Y
 4 +Z
 5 -Z
 
 */

void f_faceMesh(const uchar mask, float3 vertex[4], float3* normal){
    switch(mask){
        case 0:
            vertex[0] = (float3)(0.5, 0.5, -0.5);
            vertex[1] = (float3)(0.5, 0.5, 0.5);
            vertex[2] = (float3)(0.5, -0.5, 0.5);
            vertex[3] = (float3)(0.5, -0.5, -0.5);
            
            *normal = (float3)(1.0, 0.0, 0);
            break;
        case 1:
            vertex[0] = (float3)(-0.5, 0.5, 0.5);
            vertex[1] = (float3)(-0.5, 0.5, -0.5);
            vertex[2] = (float3)(-0.5, -0.5, -0.5);
            vertex[3] = (float3)(-0.5, -0.5, 0.5);
            
            *normal = (float3)(-1.0, 0.0, 0);
            break;
        case 2:
            vertex[0] = (float3)(-0.5, 0.5, 0.5);
            vertex[1] = (float3)(0.5, 0.5, 0.5);
            vertex[2] = (float3)(0.5, 0.5, -0.5);
            vertex[3] = (float3)(-0.5, 0.5, -0.5);
            
            *normal = (float3)(0.0, 1.0, 0);
            break;
        case 3:
            vertex[0] = (float3)(-0.5, -0.5, -0.5);
            vertex[1] = (float3)(0.5, -0.5, -0.5);
            vertex[2] = (float3)(0.5, -0.5, 0.5);
            vertex[3] = (float3)(-0.5, -0.5, 0.5);
            
            *normal = (float3)(0.0, -1.0, 0);
            break;
        case 4:
            vertex[0] = (float3)(0.5, 0.5, 0.5);
            vertex[1] = (float3)(-0.5, 0.5, 0.5);
            vertex[2] = (float3)(-0.5, -0.5, 0.5);
            vertex[3] = (float3)(0.5, -0.5, 0.5);
            
            *normal = (float3)(0.0, 0.0, 1.0);
            break;
        case 5:
            vertex[0] = (float3)(-0.5, 0.5, -0.5);
            vertex[1] = (float3)(0.5, 0.5, -0.5);
            vertex[2] = (float3)(0.5, -0.5, -0.5);
            vertex[3] = (float3)(-0.5, -0.5, -0.5);
            
            *normal = (float3)(0.0, 0.0, -1.0);
            break;
    }
}
ulong4 f_meshVertexHash(const uchar mask, const ushort3 coord){
    ulong4 plane;
    switch(mask){
        case 0:
            plane.s0 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z));
            plane.s1 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z+1));
            plane.s2 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z+1));
            plane.s3 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z));
            //normal = (float3)(1.0, 0.0, 0);
            break;
        case 1:
            plane.s0 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z+1));
            plane.s1 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z));
            plane.s2 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z));
            plane.s3 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z+1));
            //normal = (float3)(-1.0, 0.0, 0);
            break;
        case 2:
            plane.s0 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z+1));
            plane.s1 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z+1));
            plane.s2 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z));
            plane.s3 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z));
            //normal = (float3)(0.0, 1.0, 0);
            break;
        case 3:
            plane.s0 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z));
            plane.s1 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z));
            plane.s2 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z+1));
            plane.s3 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z+1));
            //normal = (float3)(0.0, -1.0, 0);
            break;
        case 4:
            plane.s0 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z+1));
            plane.s1 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z+1));
            plane.s2 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z+1));
            plane.s3 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z+1));
            //normal = (float3)(0.0, 0.0, 1.0);
            break;
        case 5:
            plane.s0 = f_getCoordHash((ushort3)(coord.x, coord.y+1, coord.z));
            plane.s1 = f_getCoordHash((ushort3)(coord.x+1, coord.y+1, coord.z));
            plane.s2 = f_getCoordHash((ushort3)(coord.x+1, coord.y, coord.z));
            plane.s3 = f_getCoordHash((ushort3)(coord.x, coord.y, coord.z));
            //normal = (float3)(0.0, 0.0, -1.0);
            break;
    }
    return plane;
}

//面片偏移
inline float4 f_faceGlobalOffset(float3 vertex, float3 offset){
    return (float4)(vertex.x+(offset.x), vertex.y+(offset.y), vertex.z+(offset.z), 1.0);
}
inline ulong f_SamplingHash_27(ushort3 coord, uchar dir) {
    /*
     z
     |
     |-----------| |-----------| |-----------|
     | 6 | 7 | 8 | |14 |15 |16 | |23 |24 |25 |
     |-----------| |-----------| |-----------|
     | 3 | 4 | 5 | |12 |.  |13 | |20 |21 |22 |
     |-----------| |-----------| |-----------|
     | 0 | 1 | 2 | | 9 |10 |11 | |17 |18 |19 |
     |-------------------------------------------x
     z
    */

    /*switch(dir){
        case 0:
            if(coord.x && coord.y && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y-1, coord.x-1));
            break;
        case 1:
            if(coord.y && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y-1, coord.x));
            break;
        case 2:
            if(coord.x+1 < XDomainMax && coord.y && coord.z) return upsample((uint)(coord.z-1), upsample((ushort)(coord.y-1), coord.x+1));
            break;
        case 3:
            if(coord.x && coord.y) return upsample((uint)(coord.z), upsample((ushort)(coord.y-1), coord.x-1));
            break;
        case 4:
            if(coord.y) return upsample((uint)(coord.z), upsample(coord.y-1, coord.x));
            break;
        case 5:
            if(coord.x+1 < XDomainMax && coord.y) return upsample((uint)(coord.z), upsample(coord.y-1, coord.x+1));
            break;
        case 6:
            if(coord.x && coord.y && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y-1, coord.x-1));
            break;
        case 7:
            if(coord.y && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y-1, coord.x));
            break;
        case 8:
            if(coord.x+1 < XDomainMax && coord.y && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y-1, coord.x+1));
            break;

        case 9:
            if(coord.x && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y, coord.x-1));
            break;
        case 10:
            if(coord.z) return upsample((uint)(coord.z-1), upsample(coord.y, coord.x));
            break;
        case 11:
            if(coord.x+1 < XDomainMax && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y-1, coord.x+1));
            break;
        case 13:
            if(coord.x+1 < XDomainMax) return upsample((uint)(coord.z), upsample(coord.y, coord.x+1));
            break;
        case 14:
            if(coord.x && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y, coord.x));
            break;
        case 15:
            if(coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y, coord.x+1));
            break;
        case 16:
            if(coord.x+1 < XDomainMax && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y, coord.x+1));
            break;

        case 17:
            if(coord.x && coord.y+1 < YDomainMax && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y+1, coord.x-1));
            break;
        case 18:
            if(coord.y+1 < YDomainMax && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y+1, coord.x));
            break;
        case 19:
            if(coord.x+1<XDomainMax && coord.y+1 < YDomainMax && coord.z) return upsample((uint)(coord.z-1), upsample(coord.y+1, coord.x+1));
            break;
        case 20:
            if(coord.x && coord.y+1 < YDomainMax) return upsample((uint)(coord.z), upsample(coord.y+1, coord.x-1));
            break;
        case 21:
            if(coord.y+1 < YDomainMax) return upsample((uint)(coord.z), upsample(coord.y+1, coord.x));
            break;
        case 22:
            if(coord.x+1<XDomainMax && coord.y+1 < YDomainMax) return upsample((uint)(coord.z), upsample(coord.y+1, coord.x+1));
            break;
        case 23:
            if(coord.x && coord.y+1 < YDomainMax && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y+1, coord.x-1));
            break;
        case 24:
            if(coord.y+1 < YDomainMax && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y+1, coord.x));
            break;
        case 25:
            if(coord.x+1<XDomainMax && coord.y+1 < YDomainMax && coord.z+1 < ZDomainMax) return upsample((uint)(coord.z+1), upsample(coord.y+1, coord.x+1));
            break;
    }*/

}
//网格组装
void f_MeshAssemble(__global float4* vertex, __global float4* normal, __global ulong2* hash, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum){
    const size_t id = CurrentID;
    if(id<eNum){
        uchar currentMask = rangeNum[id];
        if(currentMask<0x3f){
            ushort3 coord = f_getHashCoord(hash[id].s1);
            float3 newCoord = (float3)(coord.x, coord.y, coord.z);
            
            //查找27宫格

            newCoord -= 32767.5f;
            newCoord *= 0.1f;
            //coord -= (uint3)(32767, 32767, 32767);
            //coord = 0;
            uint VertexEndID = rangeOffset[id].s0;
            //if(VertexEndID >= eNum*16) return;
            
            //构建六面数据
            for(uchar i=0; i<6; ++i){
                if(!(currentMask&(1<<i))){
                    float3 face[4];
                    float3 newNormal;
                    //取得面顶点坐标
                    f_faceMesh(i, face, &newNormal);
                    
                    //面数据拷贝到顶点数组中
                    VertexEndID -= 4;
                    vertex[VertexEndID] = f_faceGlobalOffset(face[0]*0.1f, newCoord);
                    vertex[VertexEndID+1] = f_faceGlobalOffset(face[1]*0.1f, newCoord);
                    vertex[VertexEndID+2] = f_faceGlobalOffset(face[2]*0.1f, newCoord);
                    vertex[VertexEndID+3] = f_faceGlobalOffset(face[3]*0.1f, newCoord);
                    
                    ulong4 newHash = f_meshVertexHash(i, coord);
                    hash[VertexEndID  ].s0 = newHash.s0;
                    hash[VertexEndID+1].s0 = newHash.s1;
                    hash[VertexEndID+2].s0 = newHash.s2;
                    hash[VertexEndID+3].s0 = newHash.s3;
                    /*if(id==0){
                        //printf("    f_MeshAssemble %u\n", VertexEndID);
                        printf("v4 = %v4u \n", newHash);
                    }*/
                    /*rangeOffset[VertexEndID  ].s1 = VertexEndID;
                    rangeOffset[VertexEndID+1].s1 = VertexEndID+1;
                    rangeOffset[VertexEndID+2].s1 = VertexEndID+2;
                    rangeOffset[VertexEndID+3].s1 = VertexEndID+3;*/
                    
                    for(uchar normalOffset=0; normalOffset<4; ++normalOffset) normal[ VertexEndID + normalOffset ].xyz = newNormal;
                    
                }
            }
        }
    }
    //if(id==0){
    //    printf("rangeOffset = %u ", rangeOffset[eNum-1].s0);
    //    printf("62&1  %u \n", 62&1);
    //}
}

//排列网格顶点



